/*
*  This file is part of ygg-brute
*  Copyright (c) 2020 ygg-brute authors
*  See LICENSE for licensing information
*/

#include "common.h"

#define DECLSPEC DEVICE

#include "generic/const_def.h"
#include "generic/const_2526_def.h"
#include "generic/bignum.h"
#include "generic/field.h"
#include "generic/point.h"
#include "generic/random.h"
#include "generic/scalar_mul.h"

#undef DECLSPEC

#include "sha512.h"
#include "node_id.h"

__constant FieldElement ed25519_d2 = ED25519_D2;

__constant Bn256 l_3 = BN_L_3;
__constant Bn256 l_4 = BN_L_4;
__constant Bn256 l_5 = BN_L_5;
__constant Bn256 l_6 = BN_L_6;
__constant Bn256 l_7 = BN_L_7;
__constant Bn256 l_7_ub = BN_L_7_UB;
__constant Bn256 l_minus_one = BN_L_MINUS_ONE;

__constant FieldElement two = {2, 0, 0, 0, 0, 0, 0, 0, 0, 0};

#define KERNEL_SPEC(NAME, BSIZE) __kernel void NAME

#include "generic/kernel/common.h"

#include "generic/kernel/precompute_addends.h"
#include "generic/kernel/compute_midstate.h"
#include "generic/kernel/compute_uv.h"
#include "generic/kernel/batch_invert.h"

__kernel void compute_address_kernel(
    GLOBAL uint8_t* addresses,
    GLOBAL uint32_t* __restrict__ us,
    GLOBAL const uint32_t* __restrict__ vs
) {
    const size_t gid = global_id();
    const size_t gsize = global_size();

    uint32_t node_id[8];

    FieldElement t, u, v;
    fe_load_global(u, us);
    fe_load_global(v, vs);

    fe_mul(t, u, v);

    uint32_t compact_fe[10];
    fe_zero(compact_fe);

    fe_to_bytes_no_reduce((uint8_t*)(compact_fe), t);
    sha512_of_32_byte_block(compact_fe, node_id);

    uint8_t address[16];
    node_id_to_ipv6(node_id, address);

    for(size_t i = 0; i < 16; ++i) {
        addresses[16 * gid + i] = address[i];
    }
}